﻿#include "md_cl.hpp"

/////TODO: написать освобождение ресурсов распомотреть случа падении программы !!!!
int GetNumfloor(int maxalive)
{	int numParticle=0;
	int numfloor=0;
//maxalive-4 больше чем допустимое количество одновременни считающихся нитей
// запустить максимальное допустимое количество одновременни считающихся нитей
// с учетом, что последний элемент должен имень не четный индекс ( считаем димеры)
//Nramk - колчиство слоев которые могут одновременно считаться Nramk GROUP_ALIVE_MIN*24 GROUP_ALIVE_MIN 9
//	groupSize=pf*GROUP_ALIVE_MIN;
		if((maxalive-4)>Nramk)
		{
			numfloor=Nramk;
		}
		else
		{
			numfloor=2*((maxalive)/2+ maxalive%2)-4;		
		}
		numfloor = ((numfloor < GROUP_ALIVE_MIN) ? GROUP_ALIVE_MIN : numfloor);
		
		numParticle = (numfloor/GROUP_ALIVE_MIN+ ((numfloor%GROUP_ALIVE_MIN==0)?0:1))*GROUP_ALIVE_MIN;
		if (numParticle>(N-4))
		{
			numParticle=(N-4);
		}
		
		return numParticle;
	//status=clEnqueueNDRangeKernel( queue,Calc_ALL,1,	NULL,globalThreads,localThreads, 0, NULL, &myEvent2);
}
cl_int GetNumThreads(int maxalive)
{
 int numParticle=0;
 int numfloor=0;
//maxalive-4 больше чем допустимое количество одновременни считающихся нитей
// запустить максимальное допустимое количество одновременни считающихся нитей
// с учетом, что последний элемент должен имень не четный индекс ( считаем димеры)
//Nramk - колчиство слоев которые могут одновременно считаться Nramk GROUP_ALIVE_MIN*24 GROUP_ALIVE_MIN 9
//	groupSize=pf*GROUP_ALIVE_MIN;
		if((maxalive-4)>Nramk)
		{
			numfloor=Nramk;
		}
		else
		{
			numfloor=2*((maxalive)/2+ maxalive%2)-4;		
		}
		numfloor = ((numfloor < GROUP_ALIVE_MIN) ? GROUP_ALIVE_MIN : numfloor);
		
		numParticle = (numfloor/GROUP_ALIVE_MIN+ ((numfloor%GROUP_ALIVE_MIN==0)?0:1))*GROUP_ALIVE_MIN*pf;
		// если окажется что нужное количество нитий превосходит N  тогда возьмем столько сколько надо.
		if (numParticle>(N-4)*pf)
		{
			numParticle=(N-4)*pf;

		}
	//TODO: Всегда все считаем!!!!
	numParticle=(N-4)*pf;
 return (cl_int)numParticle;
}

cl_int GetGroupSize()
{
return GROUP_ALIVE_MIN*pf;
}
void Set_N_shift()
{
}


void doDef(float *x,float *z,float *t)
	{
	float	yty=0.0;
	float	yty_1=0.0f;
	int index_2;
	int index_2_1;
	for (int i=0; i<13;i++)
		{
		yty=0.0;
		yty_1=0.0f;
		for (int j =0; j<N; j++)
			{	index_2=(i+j*pf);
		index_2_1=(i+(j-1)*pf);
		if(j<N-8)
			{

			x[index_2]=V;
			z[index_2]=4*j+h*i;
			t[index_2]=0.0f;

			}
		else
			{yty+=T0;

		t[index_2]=yty;
		z[index_2]=z[index_2_1]+r*cos(yty)+r*cos(yty_1);
		x[index_2]=x[index_2_1]+r*sin(yty)+r*sin(yty_1);

		yty_1+=T0;
			}
			}
		}
	}
void doDef_rost(float *x,float *z,float *t)
	{
	float	yty=0.0;
	float	yty_1=0.0f;
	int index_2;
	int index_2_1;
	for (int i=0; i<13;i++)
		{
		yty=0.0;
		yty_1=0.0f;
		for (int j =0; j<(N/2); j++)
			{	index_2=(i+j*pf);
		index_2_1=(i+(j-1)*pf);
		if(j<(N/2)-8)
			{

			x[index_2]=V;
			z[index_2]=4*j+h*i;
			t[index_2]=0.0f;

			}
		else
			{yty+=T0;

		t[index_2]=yty;
		z[index_2]=z[index_2_1]+r*cos(yty)+r*cos(yty_1);
		x[index_2]=x[index_2_1]+r*sin(yty)+r*sin(yty_1);

		yty_1+=T0;
			}
			}
		}
	}

std::string convertToString(const char *filename)
	{
	size_t size;
	char*  str;
	std::string s;

	std::fstream f(filename, (std::fstream::in | std::fstream::binary));

	if(f.is_open())
		{
		size_t fileSize;
		f.seekg(0, std::fstream::end);
		size = fileSize = f.tellg();
		f.seekg(0, std::fstream::beg);

		str = new char[size+1];
		if(!str)
			{
			f.close();
			return NULL;
			}
		f.read(str, fileSize);
		f.close();
		str[size] = '\0';	
		s = str;		
		return s;
		}
	return NULL;
	}

int main(int argc, char ** argv)
	{
	size_t kernelWorkGroupSize;         /**< Group size returned by kernel */
    size_t groupSize;                   /**< Work-Group size */
	
	//int numBodies;      /**< No. of particles*/
	printf("sizeof(long)=%d\n sizeof(long long int )=%d\n",sizeof(long),sizeof(long long int ));
	time_t seconds;
	time_t second0;
	float _time[1];
	_time[0]=0.0f;
	int maxalive=N;
	int minDead=N;
	cl_int numParticles=Nramk;
	
	bool isGpu=true;
	bool isATI=true;
	int devID=-1;
	for(int i = 1; i < argc; i++)
    {
        if(!strcmp("cpu", argv[i]))
            isGpu = false;
    }
	for(int i = 1; i < argc; i++)
    {
        if(!strcmp("NV", argv[i]))
            isATI = false;
    }
	for(int i = 1; i < argc; i++)
    {
        if(!strcmp("D0", argv[i]))
            devID = 0;
			if(!strcmp("D1", argv[i]))
            devID = 1;
			if(!strcmp("D2", argv[i]))
            devID = 2;
			if(!strcmp("D3", argv[i]))
            devID = 3;
			if(!strcmp("D4", argv[i]))
            devID = 4;
			if(!strcmp("D5", argv[i]))
            devID = 5;
			if(!strcmp("D6", argv[i]))
            devID = 6;
			if(!strcmp("D7", argv[i]))
            devID = 7;
    }
	printf("Device %d\n",devID);
	
	
	float host_B_2__= 1.40f;//0.315//LONG
	float host_A_1___= 0.7f;//0,421319276064 = 100KT//LATER
	float host_K_2___= 1.2850237919952f;//2.0f*292.0f*305.16f*1.3806504e-5f//2,0f*292,0f*305,16f*1,3806504e-5f
	float host_T0= 0.2f; 
	float host_T1= 0.0f; 
	float host_rho= 3.52f; //long
	float host_rh1= 3.90f;//later 
	float host_g1= 3.0e-9f;//  eta=0.004   3.0e-10
	float host_g2= 6.5e-8f;//  6.5e-9
	float host_dt= 1.0e-10f;
	float host_r0= 0.2f;
	float host_r1= 0.2f;
	//float host_Ot=0.04f;
	//float host_Ot_long=0.04f;
	float host_Ot=0.0f;
	float host_Ot_long=0.0f;

	float host_koe= 0.1f;
	float host_koe_long= 0.1f;
	float host_dt_g1_1=host_dt/host_g1;// 0.3333333333e-1f;
	float host_dt_g2_1= host_dt/host_g2;//0.1538461538e-2f;
	float host_T=305.16f;//305.16f //305.16f 
	float host_K= 1.3806504e-5f; // 10^9*10^9
	int NWITEMS= N*pf;
	long N_w=1000; //Кол-во запусков без вывода
	groupSize=pf*GROUP_ALIVE_MIN;
	
	
    size_t maxWorkGroupSize;            /**< Max allowed work-items in a group */
    cl_uint maxDimensions;              /**< Max group dimensions allowed */
    size_t* maxWorkItemSizes;           /**< Max work-items sizes in each dimensions */

	cl_float3 * host_cc=(cl_float3*)malloc(pf*N*sizeof(cl_float3));	
	cl_float3 * host_lr=(cl_float3*)malloc(2*pf*N*sizeof(cl_float3));
	cl_float3 * host_ud=(cl_float3*)malloc(2*pf*N*sizeof(cl_float3));

	cl_float * host_tr=(cl_float*)malloc(pf*N*sizeof(cl_float));
	cl_float * host_tdis=(cl_float*)malloc(pf*N*sizeof(cl_float));
	cl_float * host_tdisup=(cl_float*)malloc(pf*N*sizeof(cl_float));
	cl_float * host_inuse=(cl_float*)malloc(pf*N*sizeof(cl_float));
	cl_ulong * host_mw=(cl_ulong*)malloc(pf*N*sizeof(cl_ulong));	
	bool  done__[pf*N];
	bool  done__up[pf*N];
	bool  done__2[pf];
	cl_uint host_N_shift[13];
	cl_float  host_x[pf*N];	
	cl_float  host_z[pf*N];
	cl_float  host_t[pf*N];
	
	bool spfind =false;
	bool spfirstfind =false;
	unsigned int curpos6=N;
	unsigned int  curpos13=N-8;
	unsigned int   SPEED__[N];
	float ZMAX=0.0f;
	cl_ulong startTime, endTime;
		startTime=0; endTime=0;	
	FILE *file;
	FILE *file2;
	FILE *file3;
	FILE *file4;
	FILE *file5;
	FILE *file6;
	FILE *file7;
	FILE *file8;
	FILE *file9;
	FILE *file_cl;
	FILE *file_cl_done;
	char buf[50];//=(char*)malloc(50*sizeof(char));
	time (&second0);
	struct tm * tblock;
	tblock = localtime(&second0);
	sprintf(buf,"%04d-%02d-%02d_%02d%02d%02d",tblock->tm_year+1900,tblock->tm_mon,tblock->tm_mday,tblock->tm_hour,tblock->tm_min,tblock->tm_sec);
	printf("Program 's begun work at %s\n",buf);
	char file_name[100];//=(char*)malloc(50*sizeof(char));
	char file_name2[100];//=(char*)malloc(50*sizeof(char));
	char file_name3[100];//=(char*)malloc(50*sizeof(char));
	char file_name4[100];//=(char*)malloc(50*sizeof(char));
	char file_name5[100];//=(char*)malloc(50*sizeof(char));
	char file_name6[100];//=(char*)malloc(50*sizeof(char));
	char file_name7[100];//=(char*)malloc(50*sizeof(char));
	char file_name8[100];
	char file_name9[100];
	char file_name_cl_done[50];//=(char*)malloc(50*sizeof(char));
	char buf_seek[64];
	char * pch;
	char  cmd_[60];
	//TODO: надо проверить что все нормально работает!
	strcpy(cmd_,"mkdir ");
	system(strcat(cmd_,buf));
	strcpy(cmd_,"cd ");
	system(strcat(cmd_,buf));


	strcpy(file_name,buf);
	strcpy(file_name2,buf);
	strcpy(file_name3,buf);
	strcpy(file_name4,buf);
	strcpy(file_name5,buf);
	strcpy(file_name6,buf);
	strcpy(file_name7,buf);
	strcpy(file_name8,buf);
	strcpy(file_name9,buf);
	strcat(file_name,"/");
	strcat(file_name2,"/");
	strcat(file_name3,"/");
	strcat(file_name4,"/");
	strcat(file_name5,"/");
	strcat(file_name6,"/");
	strcat(file_name7,"/");
	strcat(file_name8,"/");
	strcat(file_name9,"/");
	strcat(file_name,"!file_");
	strcat(file_name,buf);
	strcat(file_name,".txt");
	strcat(file_name2,"!pot_");
	strcat(file_name2,buf);
	strcat(file_name2,".txt");
	strcat(file_name3,"!param_");
	strcat(file_name3,buf);
	strcat(file_name3,".txt");
	strcat(file_name4,"!NXYTudrl_");
	strcat(file_name4,buf);
	strcat(file_name4,".bin");
	strcat(file_name5,"!timing_");
	strcat(file_name5,buf);
	strcat(file_name5,".txt");
	strcat(file_name6,"!random_");
	strcat(file_name6,buf);
	strcat(file_name6,".txt");
	strcat(file_name7,"!uptiming_");
	strcat(file_name7,buf);
	strcat(file_name7,".txt");
	strcat(file_name8,"!dinum_");
	strcat(file_name8,buf);
	strcat(file_name8,".txt");
	strcat(file_name9,"!speed_");
	strcat(file_name9,buf);
	strcat(file_name9,".txt");
	
	strcpy(file_name_cl_done,buf);
	strcat(file_name_cl_done,"/");

	strcat(file_name_cl_done,"!cl_");
	strcat(file_name_cl_done,buf);
	strcat(file_name_cl_done,".cl");
	file=	fopen( file_name, "w" );
	file2=	fopen( file_name2, "w" );
	file3=	fopen( file_name3, "w" );
	file4=	fopen( file_name4, "wb" );
	file5=	fopen( file_name5, "w" );
	file6=	fopen( file_name6, "w" );
	file7=	fopen( file_name7, "w" );
	file8=	fopen( file_name8, "w" );
	file9=	fopen( file_name9, "w" );
	//file_cl=fopen("1.cl","r");
	file_cl_done=fopen(file_name_cl_done,"w");
	if(file==NULL || file2==NULL|| file3==NULL||file4==NULL||file5==NULL||file6==NULL||file7==NULL||file8==NULL||file9==NULL)
	{
		printf("File isn`t open!");
		getchar();
		return 1;
	}
		
	printf("1 Unit of length = 1 nm\n");
	printf("1 Unit of time = 1 second\n");
	printf("1 Unit of temperature = 1 Kelvin\n");
	printf("K bolcman: K=%e\n",host_K);
	printf("N items %d x 13 : N=%d, Nramk=%d, GROUP_ALIVE_MIN=%d \n",N,NWITEMS,Nramk,GROUP_ALIVE_MIN);
	printf("Angle def GDP: T0=%e\n",host_T0);
	printf("Angle def GTP: T1=%e\n",host_T1);	
	printf("Viscosity shift: g1=%e\n",host_g1);
	printf("Viscosity rotate: g2=%e\n",host_g2);
	printf("Step time dt=%e\n",host_dt);
	printf("Temperature T=%e K\n",host_T);
	printf("#KT=%e\n",(host_K*host_T));
	printf("host_rho=%e\n",host_rho);
	printf("host_rh1=%e\n",host_rh1);
	printf("host_koe=%e host_koe_long=%e\n",host_koe,host_koe_long);
	printf("host_Ot=%e host_Ot_long=%e\n",host_Ot,host_Ot_long);
	printf("Constant of angel potencial        K_2___=%e; K_2___= %e KT\n",host_K_2___,(host_K_2___/(host_K*host_T)));
	printf("Constant of lateral potencial      A_1___=%e; A_1___= %e KT\n",host_A_1___,(host_A_1___/(host_K*host_T)));
	printf("Constant of lateral potencial r1=%e nm\n",host_r1);
	printf("Constant of longitudinal potencial B_2__ =%e; B_2__ = %e KT \n",host_B_2__,(host_B_2__/(host_K*host_T)));
	printf("Constant of longitudinal potencial r0=%e nm\n",host_r0);
	printf("N step NpFrame = %d,N_w=%d ,N_OUT_ALL= %d,  N_all= %e\n",NfoPr,N_w,N_OUT_ALL,(float)N_w*(float)N_OUT_ALL*(float)NfoPr);
	
		
	fprintf(file_cl_done,"__constant float K=%ef;\n",host_K);
	fprintf(file_cl_done,"__constant float rho=%ef;\n",host_rho);
	fprintf(file_cl_done,"__constant float rh1=%ef;\n",host_rh1);
	fprintf(file_cl_done,"__constant float koe=%ef;\n",host_koe);
	fprintf(file_cl_done,"__constant float Ot_long=%ef;\n",host_Ot_long);
	fprintf(file_cl_done,"__constant float koe_long=%ef;\n",host_koe_long);
	fprintf(file_cl_done,"__constant float Ot=%ef;\n",host_Ot);
	fprintf(file_cl_done,"__constant float r0=%ef;\n",host_r0);
	fprintf(file_cl_done,"__constant float r1=%ef;\n",host_r1);
	
	fprintf(file_cl_done,"__constant float _r04=%ef;\n",2.5f);//важные параметры обрезания!!!!
	fprintf(file_cl_done,"__constant float _r12=%ef;\n",2.5f);
	
	fprintf(file_cl_done,"__constant float _1r0=%ef;\n",1.0f/host_r0);
	fprintf(file_cl_done,"__constant float _1r1=%ef;\n",1.0f/host_r1);

	fprintf(file_cl_done,"__constant float T0=%ef;\n",host_T0);
	fprintf(file_cl_done,"__constant float T1=%ef;\n",host_T1);
	fprintf(file_cl_done,"__constant float g1=%ef;\n",host_g1);
	fprintf(file_cl_done,"__constant float g2=%ef;\n",host_g2);
	fprintf(file_cl_done,"__constant float dt_g1_1=%ef;\n",host_dt_g1_1);
	fprintf(file_cl_done,"__constant float dt_g2_1=%ef;\n",host_dt_g2_1);
	fprintf(file_cl_done,"__constant float dt=%ef;\n",host_dt);
	fprintf(file_cl_done,"__constant float T=%ef;\n",host_T);
	fprintf(file_cl_done,"__constant float K_2___=%ef;\n",host_K_2___);
	fprintf(file_cl_done,"__constant float A_1___=%ef;\n",host_A_1___);
	fprintf(file_cl_done,"__constant float B_2__=%ef;\n",host_B_2__);	
	fprintf(file_cl_done,"__constant int NfoPr = %d;\n",NfoPr);
	

	fprintf(file3,"1 Unit of length = 1 nm\n");
	fprintf(file3,"1 Unit of time = 1 second\n");
	fprintf(file3,"1 Unit of temperature = 1 Kelvin\n");
	fprintf(file3,"K bolcman: K=%e\n",host_K);
	fprintf(file3,"N items %d x 13 : N=%d, Nramk=%d, GROUP_ALIVE_MIN=%d \n",N,NWITEMS,Nramk,GROUP_ALIVE_MIN);
	//fprintf(file3,"N items %d x 13 : N=%d\n",N,NWITEMS);
	fprintf(file3,"Angle def GDP: T0=%e\n",host_T0);
	fprintf(file3,"Angle def GTP: T1=%e\n",host_T1);	
	fprintf(file3,"Viscosity shift: g1=%e\n",host_g1);
	fprintf(file3,"Viscosity rotate: g2=%e\n",host_g2);
	fprintf(file3,"Step time dt=%e\n",host_dt);
	fprintf(file3,"Temperature T=%e K\n",host_T);
	fprintf(file3,"KT=%e\n",(host_K*host_T));	
	fprintf(file3,"host_rho=%e\n",host_rho);
	fprintf(file3,"host_rh1=%e\n",host_rh1);
	fprintf(file3,"host_koe=%e host_koe_long=%e\n",host_koe,host_koe_long);
	fprintf(file3,"host_Ot=%e host_Ot_long=%e\n",host_Ot,host_Ot_long);
	fprintf(file3,"Constant of angel potencial        K_2___=%e; K_2___= %e KT\n",host_K_2___,(host_K_2___/(host_K*host_T)));
	fprintf(file3,"Constant of lateral potencial      A_1___=%e; A_1___= %e KT\n",host_A_1___,(host_A_1___/(host_K*host_T)));
	fprintf(file3,"Constant of lateral potencial r1=%e nm\n",host_r1);
	fprintf(file3,"Constant of longitudinal potencial B_2__ =%e; B_2__ = %e KT \n",host_B_2__,(host_B_2__/(host_K*host_T)));
	fprintf(file3,"Constant of longitudinal potencial r0=%e nm\n",host_r0);
	fprintf(file3,"N step NpFrame = %d,N_w=%d ,N_OUT_ALL= %d,  N_all= %e\n",NfoPr,N_w,N_OUT_ALL,(float)N_OUT_ALL*(float)NfoPr);
	fprintf(file3,"Program 's begun work at %s\n",buf);
	
	fclose(file3);	
	fclose(file_cl_done);
	file_cl=fopen("1.cl","rb");
	file_cl_done=fopen(file_name_cl_done,"ab");
	const int buf_size = 15638;
	fseek(file_cl_done,0,SEEK_END);
	fseek(file_cl,0,SEEK_END);

       long int src_size = ftell(file_cl);
    fseek(file_cl,0,SEEK_SET);
 
        // Количество целых проходов (для частей файлы, которые полностью умещаются в буфер)
        size_t a_mount = src_size / buf_size;
        // Остаток (остаток файла)
        size_t b_mount = src_size % buf_size;
 
        
        // Создание потока для файла-копии
       
        char buf_2[buf_size];
 
        // Цикл по числу полных проходов
        for(size_t i = 0; i < a_mount; ++i)
        {
		fread (buf_2,1,buf_size,file_cl);
		fwrite (buf_2,1,buf_size,file_cl_done);
        }
        
        // Если есть остаток
        if(b_mount != 0)
        {
		fread (buf_2,1,b_mount,file_cl);
		fwrite (buf_2,1,b_mount,file_cl_done);
        }
 
        fclose(file_cl);
        fclose(file_cl_done);
 
        
	seconds = time (NULL);
	printf("Time step #: %d , seconds from begin: %ld\n",0,(seconds-second0));
int status;

	for(int i=0;i<N*pf;i++)
	{
		host_tr[i]=host_T0;
		host_inuse[i]=1.0f;
		host_tdis[i]=0.0f;
		host_tdisup[i]=0.0f;
		done__[i]=true;
		done__up[i]=true;
		host_mw[i]=0;
	}
	
	//cl_device_id device;
	size_t global_work_size = GROUP_ALIVE_MIN;
	size_t global_work_size_full =NWITEMS ;
	
	cl_uint platformCount;
	status=clGetPlatformIDs( 0, 0, &platformCount );
	if(status != CL_SUCCESS) 
	{ 
		printf("Error:  \
			   (clGetPlatformIDs).Error N: %d\n",status );
		return -1;
	}
	else
	{
		printf("SUCCESS: clGetPlatformIDs!\n");
	}
	cl_platform_id *platformList = new cl_platform_id[ platformCount ];
	status=clGetPlatformIDs( platformCount, platformList, 0 );
	if(status != CL_SUCCESS) 
	{ 
		printf("Error:  \
			   (clGetPlatformIDs).Error N: %d\n",status );
		return -1;
	}
	else
	{
		printf("SUCCESS: clGetPlatformIDs!\n");
	}
	size_t paramSize;
	char *platformName;
	cl_uint platformid=-1;
	for(int i=0;i<platformCount;i++)
	{

		clGetPlatformInfo( platformList[i], CL_PLATFORM_NAME, 0, 0, &paramSize );
		platformName = new char[paramSize];
		clGetPlatformInfo( platformList[i], CL_PLATFORM_NAME, paramSize, platformName, 0 );
		printf("Available platform :%d: %s\n",i,platformName);
		if(isATI)
		{
			if(!strcmp("AMD Accelerated Parallel Processing", platformName))
			   platformid=i;
		}
		else
		{
		if(!strcmp("NVIDIA CUDA", platformName))
			   platformid=i;
		}	
	}
	
	printf("Platform %d selected\n",platformid);
	if(platformid<0)
	{
	printf("Can't  find needed platform\n");
	return -1;
	}
	unsigned int deviceCount;
	cl_device_id *deviceList;
	if(isGpu)
	{	
		printf("GPU selected\n");
		clGetDeviceIDs( platformList[platformid], CL_DEVICE_TYPE_GPU, 0, 0,(cl_uint *) &deviceCount );
		printf("Finded %d device\n",deviceCount);
		if(deviceCount>0)
		{
		deviceList = new cl_device_id[ deviceCount ];
		//printf("Finded %d device\n",deviceCount);
		status=clGetDeviceIDs( platformList[platformid], CL_DEVICE_TYPE_GPU, (cl_uint)deviceCount, deviceList, NULL );
		}
		else
		{
		return -1;
		}
		
	
//	status=clGetDeviceIDs( platformList[platformid], CL_DEVICE_TYPE_GPU,1,&device,NULL);
	}
	else
	{
		printf("CPU selected\n");
		clGetDeviceIDs( platformList[platformid], CL_DEVICE_TYPE_CPU, 0, 0,(cl_uint *) &deviceCount );
		deviceList = new cl_device_id[ deviceCount ];
		printf("Finded %d device\n",deviceCount);
		status=clGetDeviceIDs( platformList[platformid], CL_DEVICE_TYPE_CPU, (cl_uint)deviceCount, deviceList, NULL );
	
	 //status=clGetDeviceIDs( platformList[platformid], CL_DEVICE_TYPE_CPU,1,&device,NULL);
	}
	
	if(status != CL_SUCCESS) 
	{ 
		printf("Error:  \
			   (clGetDeviceIDs).Error N: %d\n",status );
		return -1;
	}
	//cl_bool isAV[deviceCount];
	cl_bool isAV_t=CL_FALSE;
	for(int i=0;i<deviceCount;i++)
	{

		clGetDeviceInfo( deviceList[i], CL_DEVICE_NAME, 0, 0, &paramSize );
		platformName = new char[paramSize];
		clGetDeviceInfo( deviceList[i], CL_DEVICE_NAME, paramSize, platformName, 0 );
	
		printf("DEV:%d: %s \n",i,platformName);
	}
	
	if(devID>=0)
	{

	}
	else
	{
		
		}*/
		if(deviceCount>=0)
		{
		devID=0;
		}
	}
	if(devID<0)
	{
	printf("Erorr: Can't select good device!\n");
		return -1;
		
	}
	else
	{
	printf("Device : %d selected.\n",devID);
	
	}
	if(status != CL_SUCCESS) 
		{ 
		printf("Error:  \
			   (clGetDeviceIDs).Error N: %d\n",status );
		return -1;
		}
	else
	{printf("SUCCESS: clGetDeviceIDs!\n");
	}

	cl_context context = clCreateContext( NULL,	1,&(deviceList[devID]),0, 0, &status);
	if(status != CL_SUCCESS) 
		{ 
		printf("Error:  \
			   (clCreateContext).Error N: %d\n",status );
		return -1;
		}
	else
	{printf("SUCCESS: clCreateContext!\n");
	}
	//status=clGetContextInfo(
	cl_command_queue queue = clCreateCommandQueue( context,	deviceList[devID],	CL_QUEUE_PROFILING_ENABLE,  &status );
	if(status != CL_SUCCESS) 
	{ 
		printf("Error:  \
			   (clCreateCommandQueue).Error N: %d\n",status );
		return -1;
	}
	else
	{
		printf("SUCCESS: clCreateCommandQueue!\n");
	}

	//const char * filename  = "1.cl";
	std::string  sourceStr = convertToString(file_name_cl_done);
	const char * source    = sourceStr.c_str();
	size_t sourceSize[]    = { strlen(source) };
	
	cl_program program = clCreateProgramWithSource(
		context, 
		1, 
		&source,
		sourceSize,
		&status);
	if(status != CL_SUCCESS) 
		{ 
		printf("Error: Loading source into cl_program \
			   (clCreateProgramWithSource).Error N: %d\n",status );
		return -1;
		}
	else
	{
		printf("SUCCESS: LOADED PROGRAM!\n");
	}

	cl_int *host_init_random;

	//cl_program program = clCreateProgramWithSource( context,1,&source,0,0 ); "-cl-fast-relaxed-math -cl-mad-enable"
	status=clBuildProgram( program, 1, &(deviceList[devID]), "-cl-mad-enable", NULL, NULL );
	if(status != CL_SUCCESS) {
	printf("Error: Building Binary into cl_program (clBuildProgram).Error N: %d\n",status );
	
	}
	else{
	printf("SUCCESS:BUILD PROGRAM!\n");
	}
	
	int LOG_S=0;
	status = clGetProgramBuildInfo(program, deviceList[devID], CL_PROGRAM_BUILD_LOG, NULL, NULL, (size_t*)&LOG_S);
	if(status != CL_SUCCESS) {
	printf("Error: (clGetProgramBuildInfo).Error N: %d\n",status );
	}
   //print kernel compilation error
	   if(LOG_S>0){
	  // char programLog[LOG_S];
	   char * programLog =(char*)malloc(LOG_S*sizeof(char));


	   status = clGetProgramBuildInfo(program, deviceList[devID], CL_PROGRAM_BUILD_LOG, LOG_S, programLog, 0);
	if(status != CL_SUCCESS) {
	printf("Error: (clGetProgramBuildInfo).Error N: %d\n",status );
	}
            std::cout<<programLog<<std::endl;

           free(programLog); 
	}	
		
	cl_kernel InitRandom = clCreateKernel( program, "InitRandom", &status );
	if(status != CL_SUCCESS)
	{ printf("Error: Loading Binary into cl_program (InitRandom).Error N: %d\n",status );
	//free(programLog);
	return -1;
	}

	cl_kernel XZT_C = clCreateKernel( program, "XZT_C", &status );	
	if(status != CL_SUCCESS) 
	{printf("Error: Loading Binary into cl_program (XZT_C).Error N: %d\n",status );
	return -1;
	}
	
	cl_kernel NewKernel_KOOR = clCreateKernel( program, "NewKernel_KOOR", &status );	
	if(status != CL_SUCCESS) 
	{printf("Error: Loading Binary into cl_program (NewKernel_KOOR).Error N: %d\n",status );
	return -1;
	}
	
	cl_kernel Calc_ALL = clCreateKernel( program, "Calc_ALL", &status );	
	if(status != CL_SUCCESS)
	{printf("Error: Loading Binary into cl_program (Calc_ALL).Error N: %d\n",status );
	return -1;}
	
	status = clGetKernelWorkGroupInfo(Calc_ALL,
        deviceList[devID],
        CL_KERNEL_WORK_GROUP_SIZE,
        sizeof(size_t),
        &kernelWorkGroupSize,
        0);
	if(status != CL_SUCCESS)
	{printf("Error:clGetKernelWorkGroupInfo CL_KERNEL_COMPILE_WORK_GROUP_SIZE failed. Error N: %d\n",status );
	return -1;}
	
    if(groupSize > kernelWorkGroupSize)
    {
       
            std::cout << "Out of Resources!" << std::endl;
            std::cout << "Group Size specified : " << groupSize << std::endl;
            std::cout << "Max Group Size supported on the kernel : "
                      << kernelWorkGroupSize << std::endl;
            std::cout << "Falling back to " << kernelWorkGroupSize << std::endl;
        
			return -1;
    }else
	{
	std::cout << "Max Group Size supported on the kernel : "
                      << kernelWorkGroupSize << std::endl;
	}
 /* Get Device specific Information */
    status = clGetDeviceInfo(
        deviceList[devID],
        CL_DEVICE_MAX_WORK_GROUP_SIZE,
        sizeof(size_t),
        (void*)&maxWorkGroupSize,
        NULL);
if(status != CL_SUCCESS)
	{printf("Error:clGetDeviceInfo CL_DEVICE_MAX_WORK_GROUP_SIZE failed. Error N: %d\n",status );
	return -1;}
    
    status = clGetDeviceInfo(
        deviceList[devID],
        CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,
        sizeof(cl_uint),
        (void*)&maxDimensions,
        NULL);
if(status != CL_SUCCESS)
	{printf("Error:clGetDeviceInfo CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS failed. Error N: %d\n",status );
	return -1;}
   

    maxWorkItemSizes = (size_t*)malloc(maxDimensions * sizeof(size_t));

    status = clGetDeviceInfo(
        deviceList[devID],
        CL_DEVICE_MAX_WORK_ITEM_SIZES,
        sizeof(size_t) * maxDimensions,
        (void*)maxWorkItemSizes,
        NULL);
	if(status != CL_SUCCESS)
	{printf("Error:clGetDeviceInfo CL_DEVICE_MAX_WORK_ITEM_SIZES failed. Error N: %d\n",status );
	return -1;}
	

		numParticles=GetNumThreads(maxalive);
		printf("<<>>n=%d M:%d\n",numParticles,maxalive);
				//printf("QUQUQU22");
	for(int gf=0;gf<13;gf++)
	{
		if((maxalive-4)>Nramk)
		{
				
					host_N_shift[gf]=maxalive-Nramk;
					printf("Cur shift= %d  for pf= %d\n",host_N_shift[gf],gf);
				
				//printf("QUQUQ2wU");
		}
		else
		{
				printf("\n<<<<<<<<<<<<<>>>>>>>>>>>>>>>>>>>>>>>\n");
					host_N_shift[gf]=4 ;
					printf("Cur shift= %d  for pf= %d\n",host_N_shift[gf],gf);
		}
	//TODO: URENT затычка
		host_N_shift[gf]=4 ;
	}
 
	size_t globalThreads[] = {numParticles};
    size_t localThreads[] = {groupSize};   
	if(localThreads[0] > maxWorkItemSizes[0] ||
       localThreads[0] > maxWorkGroupSize)
    {
        std::cout << "Unsupported: Device"
            "does not support requested number of work items.";
        return -1;
    }
	
	size_t sizeInBytes = NWITEMS * sizeof(cl_int);
	host_init_random = (cl_int *) malloc(sizeInBytes);
	srand ((unsigned int) time(NULL) );
	for(int i = 0; i < NWITEMS; i++)
		{
		host_init_random[i] = cl_int(rand());
		}	
	printf("\n");

	cl_mem CL_init_random		= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_int) ,host_init_random, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_init_random into cl_program  \n");
	cl_mem CL_in1 = clCreateBuffer( context,CL_MEM_READ_WRITE,NWITEMS * sizeof(cl_uint),NULL, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_in1 into cl_program  \n");
	cl_mem CL_in2 = clCreateBuffer( context,CL_MEM_READ_WRITE,NWITEMS * sizeof(cl_uint),NULL, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_in2 into cl_program  \n");
	cl_mem CL_in3 = clCreateBuffer( context,CL_MEM_READ_WRITE,NWITEMS * sizeof(cl_uint),NULL, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_in3 into cl_program  \n");
	cl_mem CL_in4 = clCreateBuffer( context,CL_MEM_READ_WRITE,NWITEMS * sizeof(cl_uint),NULL, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_in4 into cl_program  \n");
	cl_mem CL_in5 = clCreateBuffer( context,CL_MEM_READ_WRITE,NWITEMS * sizeof(cl_uint),NULL, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_in5 into cl_program  \n");
	cl_mem buffer_out = clCreateBuffer( context,CL_MEM_READ_WRITE,NWITEMS * sizeof(cl_float),NULL, &status );
	if(status != CL_SUCCESS) printf("Error: Loading buffer_out into cl_program  \n");

		cl_mem CL_cc = clCreateBuffer( context,CL_MEM_READ_WRITE,NWITEMS * sizeof(cl_float3),NULL, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_ccx into cl_program  \n");
	cl_mem CL_lr = clCreateBuffer( context,CL_MEM_READ_WRITE,2*NWITEMS * sizeof(cl_float3),NULL, &status );
	cl_mem CL_ud = clCreateBuffer( context,CL_MEM_READ_WRITE,2*NWITEMS * sizeof(cl_float3),NULL, &status );
		cl_mem CL_tdis = clCreateBuffer( context,CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_float),host_tdis, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_tdis into cl_program  \n");
	
		cl_mem CL_tdisup = clCreateBuffer( context,CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_float),host_tdisup, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_tdisup into cl_program  \n");
	cl_mem CL_inuse = clCreateBuffer( context,CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_float),host_inuse, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_inuse into cl_program  \n");
	cl_mem CL_mw = clCreateBuffer( context,CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_ulong),host_mw, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_mw into cl_program  \n");
    
    cl_mem CL_tr = clCreateBuffer( context,CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_float),host_tr, &status );
	if(status != CL_SUCCESS) 
	{printf("Error: Loading CL_tr into cl_program \n");
	return -1;}

	float _isup[N];
	cl_uint _right__[pf*	N];
	cl_uint _left__[pf*	N];
	cl_uint _up__[pf*	N];
	cl_uint _down__[pf*	N];
	cl_uint _uT__[pf*	N];
	cl_uint _dT__[pf*	N];
	cl_uint _lT__[pf*	N];
	cl_uint _rT__[pf*	N];
	

	int pitch= pf;

	for(int i=0;i<N;i++)
	{
		_isup[i]=1.0f;
		SPEED__[i]=13;
	}
	_isup[N-1]=0.0f;
		doDef(host_x,host_z,host_t);
	//doDef_rost(host_x,host_z,host_t);
	for(int i=0;i<N;i++)
	{
		_isup[i]=1.0f;
	}
	_isup[N-1]=0.0f;
	/*for(int i=(N-1);i>(N/2);i--)
	{
	_isup[i]=0.0f;
	}
	*/
	for (int i=0;i<pf;i++)
	{
		//_isup[i+pitch *(N-1)]=0.0f;
		//host_N_shift[i]=4;
		for(int j=0;j<N;j++)
		{
			_up__[i+pitch *j]=2*(i+pitch*(j))+1;
			_down__[i+pitch *j]=2*(i+pitch *(j));
			_right__[i+pitch *j]=2*(i+pitch *j);
			_left__[i+pitch *j]=2*(i+pitch *j)+1;
			_uT__[i+pitch *j]=(i+pitch *(j));
			_dT__[i+pitch *j]=(i+pitch *(j));
			_rT__[i+pitch *j]=(i+pitch *(j));
			_lT__[i+pitch *j]=(i+pitch *(j));

		}
	}
	for (int i=1;i<pf-1;i++)
	{
		for(int j=1;j<N-1;j++)
		{
			_up__[i+pitch *j]=2*(i+pitch*(j-1));
			_down__[i+pitch *j]=2*(i+pitch *(j+1))+1;
			_right__[i+pitch *j]=2*(i+1+pitch *j)+1;
			_left__[i+pitch *j]=2*(i-1+pitch *j);
			_uT__[i+pitch *j]=(i+pitch *(j-1));
			_dT__[i+pitch *j]=(i+pitch *(j+1));
			_rT__[i+pitch *j]=(i+1+pitch *(j));
			_lT__[i+pitch *j]=(i-1+pitch *(j));
		}
	}

	for (int i=1;i<12;i++)
	{
		_up__[i+pitch *0]=2*(i+pitch*(0))+1;
		_down__[i+pitch *0]=2*(i+pitch *(0+1))+1;
		_right__[i+pitch *0]=2*(i+1+pitch *0)+1;
		_left__[i+pitch *0]=2*(i-1+pitch *0);

		_rT__[i+pitch *0]=(i+1+pitch *0);
		_lT__[i+pitch *0]=(i-1+pitch *0);

		_uT__[i+pitch *0]=(i+pitch *0);
		_dT__[i+pitch *0]=(i+pitch *1);


		_up__[i+pitch *(N-1)]=2*(i+pitch*(N-1-1));
		_down__[i+pitch *(N-1)]=2*(i+pitch *(N-1));
		_right__[i+pitch *(N-1)]=2*(i+1+pitch *(N-1))+1;
		_left__[i+pitch *(N-1)]=2*(i-1+pitch *(N-1));

		_rT__[i+pitch *(N-1)]=(i+1+pitch *(N-1));
		_lT__[i+pitch *(N-1)]=(i-1+pitch *(N-1));

		_uT__[i+pitch *(N-1)]=(i+pitch *(N-1-1));
		_dT__[i+pitch *(N-1)]=(i+pitch *(N-1));
	}

	for(int j=1;j<N-1;j++)
	{
		_up__[0+pitch *j]=2*(0+pitch*(j-1));
		_down__[0+pitch *j]=2*(0+pitch *(j+1))+1;
		_uT__[0+pitch *j]=(0+pitch *(j-1));
		_dT__[0+pitch *j]=(0+pitch *(j+1));

		_up__[12+pitch *j]=2*(12+pitch*(j-1));
		_down__[12+pitch *j]=2*(12+pitch *(j+1))+1;
		_uT__[12+pitch *j]=(12+pitch *(j-1));
		_dT__[12+pitch *j]=(12+pitch *(j+1));

		if((j>=3)&&(j<(N-3)))
		{
			_right__[0+pitch *j]=2*(1+pitch *j)+1;
			_left__[0+pitch *j]=2*(12+pitch *(j-3));

			_right__[12+pitch *j]=2*(0+pitch *(j+3))+1;
			_left__[12+pitch *j]=2*(11+pitch *j);

			_rT__[0+pitch *j]=(1+pitch *j);
			_lT__[0+pitch *j]=(12+pitch *(j-3));

			_rT__[12+pitch *j]=(0+pitch *(j+3));
			_lT__[12+pitch *j]=(11+pitch *j);
		}
		else
		{
			if(j<3)
			{
				_right__[0+pitch *j]=2*(1+pitch *j)+1;
				_left__[0+pitch *j]=2*(0+pitch *(j))+1;

				_right__[12+pitch *j]=2*(0+pitch *(j+3))+1;
				_left__[12+pitch *j]=2*(11+pitch *j);

				_rT__[0+pitch *j]=(1+pitch *j);
				_lT__[0+pitch *j]=(0+pitch *(j));

				_rT__[12+pitch *j]=(0+pitch *(j+3));
				_lT__[12+pitch *j]=(11+pitch *j);
			}
			if(j>=(N-3))
			{
				_right__[0+pitch *j]=2*(1+pitch *j)+1;
				_left__[0+pitch *j]=2*(12+pitch *(j-3));

				_right__[12+pitch *j]=2*(12+pitch *j);
				_left__[12+pitch *j]=2*(11+pitch *j);

				_rT__[0+pitch *j]=(1+pitch *j);
				_lT__[0+pitch *j]=(12+pitch *(j-3));

				_rT__[12+pitch *j]=(12+pitch *j);
				_lT__[12+pitch *j]=(11+pitch *j);
			}
		}

	}

	_up__[0+pitch *0]=2*(0+pitch*(0))+1;
	_down__[0+pitch *0]=2*(0+pitch *(0+1))+1;
	_right__[0+pitch *0]=2*(0+1+pitch *0)+1;
	_left__[0+pitch *0]=2*(0+pitch *0)+1;

	_rT__[0+pitch *0]=(0+1+pitch *0);
	_lT__[0+pitch *0]=(0+pitch *0);

	_uT__[0+pitch *0]=(0+pitch *0);
	_dT__[0+pitch *0]=(0+pitch *1);


	_up__[12+pitch *(N-1)]=2*(12+pitch*(N-1-1));
	_down__[12+pitch *(N-1)]=2*(12+pitch *(N-1));
	_right__[12+pitch *(N-1)]=2*(12+pitch *(N-1));
	_left__[12+pitch *(N-1)]=2*(12-1+pitch *(N-1));

	_rT__[12+pitch *(N-1)]=(12+pitch *(N-1));
	_lT__[12+pitch *(N-1)]=(12-1+pitch *(N-1));

	_uT__[12+pitch *(N-1)]=(12+pitch *(N-1-1));
	_dT__[12+pitch *(N-1)]=(12+pitch *(N-1));


	_up__[12+pitch *(0)]=2*(12+pitch*(0));
	_down__[12+pitch *(0)]=2*(12+pitch *(0+1))+1;
	_rT__[12+pitch *(0)]=(0+pitch *(0+3));
	_lT__[12+pitch *(0)]=(12-1+pitch *(0));

	_right__[12+pitch *(0)]=2*(0+pitch *(0+3))+1;
	_left__[12+pitch *(0)]=2*(12-1+pitch *(0));

	_uT__[12+pitch *(0)]=(12+pitch *(0));
	_dT__[12+pitch *(0)]=(12+pitch *(0+1));

	_up__[0+pitch *(N-1)]=2*(0+pitch*(N-1-1));
	_down__[0+pitch *(N-1)]=2*(0+pitch *(N-1));
	_right__[0+pitch *(N-1)]=2*(1+pitch *(N-1))+1;
	_left__[0+pitch *(N-1)]=2*(12+pitch *(N-1-3));

	_rT__[0+pitch *(N-1)]=(1+pitch *(N-1));
	_lT__[0+pitch *(N-1)]=(12+pitch *(N-1-3));

	_uT__[0+pitch *(N-1)]=(0+pitch *(N-1-1));
	_dT__[0+pitch *(N-1)]=(0+pitch *(N-1));



	cl_mem CL_isup		= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,N *sizeof(float) ,_isup, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_isup into cl_program  \n");
	
	
	cl_mem CL_N_shift		= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,13 *sizeof(cl_uint) ,host_N_shift, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_N_shift into cl_program  \n");
	
	cl_mem CL_right__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_right__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_right__ into cl_program  \n");
	
	cl_mem CL_left__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_left__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_left__ into cl_program  \n");
	
	cl_mem CL_up__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_up__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_up__ into cl_program  \n");
	
	cl_mem CL_down__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_down__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_down__ into cl_program  \n");
	
	cl_mem CL_uT__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_uT__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_uT__ into cl_program  \n");
	
	cl_mem CL_dT__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_dT__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_dT__ into cl_program  \n");
	
	cl_mem CL_lT__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_lT__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_lT__ into cl_program  \n");
	
	cl_mem CL_rT__	= clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,NWITEMS *sizeof(cl_uint) ,_rT__, &status );
	if(status != CL_SUCCESS) printf("Error: Loading CL_rT__ into cl_program  \n");
	

		cl_mem CL_x = clCreateBuffer( context,CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_float),host_x, &status );
	if(status != CL_SUCCESS) printf("Error: CL_x %d\n",status);
		cl_mem CL_z = clCreateBuffer( context,CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_float),host_z, &status );
	if(status != CL_SUCCESS) printf("Error: CL_z %d\n",status);
		cl_mem CL_t = clCreateBuffer( context,CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,NWITEMS * sizeof(cl_float),host_t, &status );
	if(status != CL_SUCCESS) printf("Error: CL_t %d\n",status);


	cl_event myEvent,myEvent2;
	cl_ulong elapsedTime ;
	clSetKernelArg(InitRandom, 0, sizeof(CL_init_random), (void*) &CL_init_random);
	clSetKernelArg(InitRandom, 1, sizeof(CL_in1), (void*) &CL_in1);
	clSetKernelArg(InitRandom, 2, sizeof(CL_in2), (void*) &CL_in2);
	clSetKernelArg(InitRandom, 3, sizeof(CL_in3), (void*) &CL_in3);
	clSetKernelArg(InitRandom, 4, sizeof(CL_in4), (void*) &CL_in4);
	clSetKernelArg(InitRandom, 5, sizeof(CL_in5), (void*) &CL_in5);
	clEnqueueNDRangeKernel( queue,InitRandom,1,	NULL,&global_work_size_full,NULL, 0, NULL, &myEvent);
	clFinish( queue );
	clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &startTime, NULL);
	clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &endTime, NULL);
	elapsedTime = endTime-startTime;
	printf("InitRandom: %d  ns\n ",elapsedTime);
			clSetKernelArg(XZT_C, 0, sizeof(CL_x), (void*) &CL_x);
			clSetKernelArg(XZT_C, 1, sizeof(CL_z), (void*) &CL_z);
			clSetKernelArg(XZT_C, 2, sizeof(CL_cc), (void*) &CL_cc);
			clEnqueueNDRangeKernel( queue,XZT_C,1,	NULL,&global_work_size_full,NULL, 0, NULL, &myEvent2);	
			clFinish( queue );
			clGetEventProfilingInfo(myEvent2, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &startTime, NULL);
			clGetEventProfilingInfo(myEvent2, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &endTime, NULL);
	
			 elapsedTime = endTime-startTime;
			printf("1 XZT_C: %d ns \n",elapsedTime);
			clSetKernelArg(NewKernel_KOOR, 0, sizeof(CL_x), (void*) &CL_x);
			clSetKernelArg(NewKernel_KOOR, 1, sizeof(CL_z), (void*) &CL_z);
			clSetKernelArg(NewKernel_KOOR, 2, sizeof(CL_t), (void*) &CL_t);
			clSetKernelArg(NewKernel_KOOR, 3, sizeof(CL_lr), (void*) &CL_lr);
			clSetKernelArg(NewKernel_KOOR, 4, sizeof(CL_ud), (void*) &CL_ud);

			clEnqueueNDRangeKernel( queue,NewKernel_KOOR,1,	NULL,&global_work_size_full,NULL, 0, NULL, &myEvent2);	
			clFinish( queue );
			clGetEventProfilingInfo(myEvent2, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &startTime, NULL);
			clGetEventProfilingInfo(myEvent2, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &endTime, NULL);
	
			elapsedTime = endTime-startTime;
			printf("2 NewKernel_KOOR: %d ns \n",elapsedTime);

	status= clEnqueueReadBuffer( queue,CL_cc,	CL_TRUE,0,	NWITEMS * sizeof(cl_float3),host_cc, NULL, NULL, NULL );
	if(status != CL_SUCCESS) printf("Error: Loading host_ccx into CL_ccx  \n");
	status=  clEnqueueReadBuffer( queue,CL_ud,	CL_TRUE,0,	2*NWITEMS * sizeof(cl_float3),host_ud, NULL, NULL,NULL );
	if(status != CL_SUCCESS) printf("Error: Loading BiCL_udxnary into host_udx  \n");
	status=  clEnqueueReadBuffer( queue,CL_lr,	CL_TRUE,0,	2*NWITEMS * sizeof(cl_float3),host_lr, NULL, NULL, NULL );
	if(status != CL_SUCCESS) printf("Error: Loading host_lrx into CL_lrx  \n");

	
	////INSERT RANDOM HERE
	fprintf(file4,"<FILE=%s;Frame=0;N=%d>;\n",buf,pf*N);
	fflush(file6);
	fclose(file6);
	fprintf(file,"#Frame\n");
	fprintf(file,"#MT\n");
	//fprintf(file4,"<Frame=0;>\n");
	int index=0;
	
	for(int j =0;j<13;j++)
	{
		for(int m=0;m<N;m++)
		{
			index=m*pf+j;
			fprintf(file,"%f;%f;%f;%f;0;%d;0;\n",host_cc[index].s[0],host_cc[index].s[1],host_cc[index].s[2],r,128*(m%2)+127);
			fprintf(file,"%f;%f;%f;%f;%d;%d;%d;\n",host_lr[2*index+1].s[0],host_lr[2*index+1].s[1],host_lr[2*index+1].s[2],r/10,0,255,0);
			fprintf(file,"%f;%f;%f;%f;%d;%d;%d;\n",host_lr[2*(index)].s[0],host_lr[2*(index)].s[1],host_lr[2*(index)].s[2],r/10,255,0,0);
			fprintf(file,"%f;%f;%f;%f;%d;%d;%d;\n",host_ud[2*(index)+1].s[0],host_ud[2*(index)+1].s[1],host_ud[2*(index)+1].s[2],r/10,0,0,255);
			fprintf(file,"%f;%f;%f;%f;%d;%d;%d;\n",host_ud[2*(index)].s[0],host_ud[2*(index)].s[1],host_ud[2*(index)].s[2],r/10,0,255,255);

		}
	}
//	fprintf(file4,"\n");
	fflush(file4);
	//fclose(file4);
	  rewind(file4);

	fclose(file4);
	status= clEnqueueWriteBuffer( queue,CL_inuse,	CL_TRUE,0,	NWITEMS * sizeof(cl_float),host_inuse, NULL, NULL, NULL );
	if(status != CL_SUCCESS) printf("Error: UPloading host_inuse into CL_inuse  \n");
	fflush(file);
	seconds = time (NULL);
	printf("Time step #: %d , seconds from begin: %ld\n",0,(seconds-second0));




			clSetKernelArg(Calc_ALL, 0, sizeof(CL_in1), (void*) &CL_in1);
			clSetKernelArg(Calc_ALL, 1, sizeof(CL_in2), (void*) &CL_in2);
			clSetKernelArg(Calc_ALL, 2, sizeof(CL_in3), (void*) &CL_in3);
			clSetKernelArg(Calc_ALL, 3, sizeof(CL_in4), (void*) &CL_in4);
			clSetKernelArg(Calc_ALL, 4, sizeof(CL_in5), (void*) &CL_in5);
			clSetKernelArg(Calc_ALL, 5, sizeof(CL_tdis), (void*) &CL_tdis);
			clSetKernelArg(Calc_ALL, 6, sizeof(CL_x), (void*) &CL_x);
			clSetKernelArg(Calc_ALL, 7, sizeof(CL_z), (void*) &CL_z);
			clSetKernelArg(Calc_ALL, 8, sizeof(CL_t), (void*) &CL_t);
			clSetKernelArg(Calc_ALL, 9, sizeof(CL_cc), (void*) &CL_cc);
			clSetKernelArg(Calc_ALL, 10, sizeof(CL_lr), (void*) &CL_lr);
			clSetKernelArg(Calc_ALL, 11, sizeof(CL_ud), (void*) &CL_ud);
			clSetKernelArg(Calc_ALL, 12, sizeof(CL_isup), (void*) &CL_isup);
			clSetKernelArg(Calc_ALL, 13, sizeof(CL_right__), (void*) &CL_right__);
			clSetKernelArg(Calc_ALL, 14, sizeof(CL_left__), (void*) &CL_left__);
			clSetKernelArg(Calc_ALL, 15, sizeof(CL_up__), (void*) &CL_up__);
			clSetKernelArg(Calc_ALL, 16, sizeof(CL_down__), (void*) &CL_down__);
			clSetKernelArg(Calc_ALL, 17, sizeof(CL_uT__), (void*) &CL_uT__);
			clSetKernelArg(Calc_ALL, 18, sizeof(CL_dT__), (void*) &CL_dT__);
			clSetKernelArg(Calc_ALL, 19, sizeof(CL_lT__), (void*) &CL_lT__);
			clSetKernelArg(Calc_ALL, 20, sizeof(CL_rT__), (void*) &CL_rT__);
			clSetKernelArg(Calc_ALL, 21, sizeof(CL_tdisup), (void*) &CL_tdisup);
			clSetKernelArg(Calc_ALL, 22, sizeof(CL_inuse), (void*) &CL_inuse);
			clSetKernelArg(Calc_ALL, 23, sizeof(CL_tr), (void*) &CL_tr);
			clSetKernelArg(Calc_ALL, 24, sizeof(CL_mw), (void*) &CL_mw);
			clSetKernelArg(Calc_ALL, 25, sizeof(CL_N_shift), (void*) &CL_N_shift);
			seconds = time (NULL);
			
	printf("Time step (BEFOR ANY CALLS OF CALC_ALL)#: %ld , seconds from begin: %ld\n",0,(seconds-second0));

	unsigned long  int sed2=0;
	for( unsigned long int  i=1;i<N_OUT_ALL;i++ )
	{			
			sed2 = time (NULL);
			ZMAX=0.0f;
			_time[0]=(float)NfoPr*(float)N_w*(float)(i-1)*(float)host_dt;
			printf("<<<<<<<<<<<<<<<<<<<<<  Calc_ALL 1x %d by %d\n",globalThreads[0],localThreads[0]);
			for(int mu=0;mu<N_w;mu++)
			{

				status=clEnqueueNDRangeKernel( queue,Calc_ALL,1,	NULL,globalThreads,localThreads, 0, NULL, &myEvent2);
				if(status != CL_SUCCESS) printf("Error:Calc_ALL error LA!!!%d \n",status);
				status=clFinish( queue );
				if(status != CL_SUCCESS) printf("Error:Calc_ALL error \n");
			}
			if(status != CL_SUCCESS) printf("Error:Calc_ALL error \n");
			clGetEventProfilingInfo(myEvent2, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &startTime, NULL);
			clGetEventProfilingInfo(myEvent2, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &endTime, NULL);	
			elapsedTime = endTime-startTime;
			printf("Calc_ALL: %ld: %e \n",(long)(i-1)*(long)(NfoPr*N_w),(float)elapsedTime*(float)1.0e-9f);
			seconds = time (NULL);
			
			
			status= clEnqueueReadBuffer( queue,CL_cc,	CL_TRUE,0,	NWITEMS * sizeof(cl_float3),host_cc, NULL, NULL, &myEvent );
			if(status != CL_SUCCESS) printf("Error: Loading host_ccx into CL_ccx  \n");
	
			status= clEnqueueReadBuffer( queue,CL_inuse,	CL_TRUE,0,	NWITEMS * sizeof(cl_float),host_inuse, NULL, NULL, NULL );
			if(status != CL_SUCCESS) printf("Error: Loading host_inuse into CL_inuse  \n");
			status= clEnqueueReadBuffer( queue,CL_tdis,	CL_TRUE,0,	NWITEMS * sizeof(cl_float),host_tdis, NULL, NULL, NULL );
			if(status != CL_SUCCESS) printf("Error: Loading host_tdis into CL_tdis  \n");
			
			status= clEnqueueReadBuffer( queue,CL_tdisup,	CL_TRUE,0,	NWITEMS * sizeof(cl_float),host_tdisup, NULL, NULL, NULL );
			if(status != CL_SUCCESS) printf("Error: Loading host_tdisup into CL_tdisup  \n");
			status= clEnqueueReadBuffer( queue,CL_t,	CL_TRUE,0,	NWITEMS * sizeof(cl_float),host_t, NULL, NULL, &myEvent2 );

			if(status != CL_SUCCESS) printf("Error: Loading host_t into CL_t  \n");
			clFinish( queue );
			clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &startTime, NULL);
			clGetEventProfilingInfo(myEvent2, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &endTime, NULL);
			clFinish( queue );
			elapsedTime = endTime-startTime;
			_time[0]=(float)NfoPr*((float)N_w)*(i)*(float)host_dt;
			printf("Mem operations : %ld: %e \n",(long)i*(long)(NfoPr*N_w),(float)elapsedTime*(float)1.0e-9f);
			
		
			
		//	printf("time spent on GPU kernel : %ld nanoeconds, step N=%ld \n", elapsedTime,N_w*NfoPr*i );
			fprintf(file,"###############%f\n",_time[0]);
			fprintf(file,"#Frame\n");
			fprintf(file,"#MT\n");
			fprintf(file2,"\t\t;;;;#Frame, time=%f\n",_time[0]);

		
			maxalive=0;
			for(int j =0;j<13;j++)
			{
				done__2[j]=false;
				
				for(int m=0;m<N;m++)
				{
					index=m*pf+j;
					//printf("%d;%d;%d : tdis %e, done %d ;\n",m,j,index,host_tdis[index],(int)done__[index]);
					if((host_tdis[index]!=0.0f)&&(done__[index]))
					{
						//printf("\t%e\n",host_tdis[index]);
						done__[index]=false;
						fprintf(file5,"%d;%d;%d;%e\n",m,j,index,_time[0]+host_dt*host_tdis[index]);
						printf("%d;%d;%d;%e\n",m,j,index,_time[0]+host_dt*host_tdis[index]);
					}
					
					if((host_tdisup[index]!=0.0f)&&(done__up[index]))
					{
						done__up[index]=false;
						fprintf(file7,"%d;%d;%d;%e\n",m,j,index,_time[0]+host_dt*host_tdisup[index]);
						printf("UP-DOWN:  %d;%d;%d;%e\n",m,j,index,_time[0]+host_dt*host_tdisup[index]);
						for(int tes=m+1;tes<N;tes++)
						{
							if(host_tdisup[tes*pf+j]!=0.0f)
							{
								fprintf(file8,"%d;%d;%d;%e;%d;%d\n",m,j,index,_time[0]+host_dt*host_tdisup[index],tes,tes-m);
								printf("N_dis UP-DOWNupto:  m=%d;%d;%d;tes=%d;%d\n",m,j,index,tes,tes-m);
								done__up[index]=true;
								break;
							}
						}
						if(!done__up[index])
						{
							fprintf(file8,"%d;%d;%d;%e;%d;%d\n",m,j,index,_time[0]+host_dt*host_tdisup[index],m,N-m);
							printf("N_dis UP-DOWN_upall:  %d;%d;%d;%d\n",m,j,index,N-m);
						}
						done__up[index]=false;
					}
					
					if((host_inuse[index]==0.0f)&&(!done__2[j]))
					{
						for(int qum=m;qum<N;qum++)
						{
							host_inuse[(qum*pf+j)]=0.0f;
							done__2[j]=true;
						}
				
					}
				
					if(host_inuse[index]==1.0f)
					{	

						
						if(m>maxalive)
						{
							maxalive=m+1;
						}
						if(	host_cc[index].s[2]>ZMAX)
						{
						ZMAX=host_cc[index].s[2];
						}
						fprintf(file,"%f;%f;%f;%f;0;%d;0;\n",host_cc[index].s[0],host_cc[index].s[1],host_cc[index].s[2],r,128*(m%2)+127);
						fprintf(file2,"%f;%d;%d;\n",host_t[index],m,j);
					}
					else
					{
					if(minDead>m)
						{
							minDead=m+1;
						}
					
					SPEED__[m]-=1;
						fprintf(file,"%f;%f;%f;%f;55;55;%d;128;\n",host_cc[index].s[0],host_cc[index].s[1],host_cc[index].s[2],r,128*(m%2)+127);
						fprintf(file2,"%f;%d;%d;\n",host_t[index],m,j);
					}

				
					host_mw[index]=0;
				}				
			}
			for(int m=0;m<N;m++)
			{
				if((SPEED__[m]<=6)&&(!spfind))
				{
					curpos6=m;
					spfind=true;
					
				}
				if((SPEED__[m]<13)&&(!spfirstfind))
				{	curpos13=m;
					spfirstfind=true;
					
				}
				SPEED__[m]=13;
			}
			spfirstfind=false;
			spfind=false;
			fprintf(file9,"%e;%d;%d;%f;\n",_time[0],curpos6,curpos13,ZMAX);
				/*	
				if(maxalive<GROUP_ALIVE_MIN)
				{
					maxalive=GROUP_ALIVE_MIN;
				}*/
				if(maxalive%2==1)
				{
					printf("!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!\n!!!!!!!!!!!!!!!\n>>>>>>>>>>>>>>>>>>>>>>>>\t");
				}
				printf("  MaxAlive=%d; minDead =%d \n",maxalive,minDead);
				if((maxalive-minDead)>Nramk)
				{
				printf("   !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!\n");
				printf("   !!!!!!!!!!!!!!!!!!!!!!!!!!!ALARM!!!!!!!!!!!!!!!!!!!!!!!!\n");
				printf("   !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!\n");
				printf("  MaxAlive=%d; minDead =%d, maxalive-minDead > Nramk \n",maxalive,minDead);
				return -1;
				}
				
				numParticles=GetNumThreads(maxalive);
				printf("<<>>n %d ; M:%d\n",numParticles,maxalive);
						//printf("QUQUQU22");
				for(int gf=0;gf<13;gf++)
				{
					if((maxalive-4)>Nramk)
					{
							
								host_N_shift[gf]=maxalive-Nramk;
								printf("Cur shift= %d  for pf= %d\n",host_N_shift[gf],gf);
							
							//printf("QUQUQ2wU");
					}
					else
					{
							//printf("\n<<<<<<<<<<<<<>>>>>>>>>>>>>>>>>>>>>>>\n");
								host_N_shift[gf]=4 ;
								printf("Cur shift= %d  for pf= %d\n",host_N_shift[gf],gf);
					}
				
					//TODO: URENT затычка
				host_N_shift[gf]=4 ;

				}
 
			 globalThreads[0] = numParticles;
			 localThreads[0] = groupSize;   
			
				printf("Group set: %d in group %d. Total groups =%d\n",numParticles,groupSize,numParticles/groupSize);
				
			   
				if(localThreads[0] > maxWorkItemSizes[0] ||
				   localThreads[0] > maxWorkGroupSize)
				{
					std::cout << "Unsupported: Device"
						"does not support requested number of work items.";
					return -1;
				}
			status= clEnqueueWriteBuffer( queue,CL_mw,	CL_TRUE,0,	NWITEMS * sizeof(cl_ulong),host_mw, NULL, NULL, NULL );
			if(status != CL_SUCCESS) printf("Error: UPLoading CL_mw into CL_mw  \n");
			status= clEnqueueWriteBuffer( queue,CL_N_shift,	CL_TRUE,0,	pf * sizeof(cl_uint),host_N_shift, NULL, NULL, NULL );
			if(status != CL_SUCCESS) printf("Error: UPLoading CL_N_shift into CL_N_shift  \n");
			status= clEnqueueWriteBuffer( queue,CL_inuse,	CL_TRUE,0,	NWITEMS * sizeof(cl_float),host_inuse, NULL, NULL, NULL );
			if(status != CL_SUCCESS) printf("Error: UPloading host_inuse into CL_inuse  \n");
			clFinish( queue );
			_time[0]=(float)NfoPr*(float)(N_w)*(i)*(float)host_dt;
			fflush(file);
			fflush(file9);
			fflush(file5);
			fflush(file8);
			fflush(file7);
			fflush(file2);
			seconds = time (NULL);
			printf("Time step #: %ld ( %e ) , seconds from begin: %ld\n",NfoPr*i*(N_w),_time[0],(seconds-second0));

	}




	

	

	fflush(file5);
	fclose(file);
	fclose(file2);
	fclose(file5);
	fclose(file7);
	fclose(file8);
	fclose(file9);

    clReleaseKernel(InitRandom);   
	clReleaseKernel(Calc_ALL);   
	clReleaseKernel(NewKernel_KOOR);   
	clReleaseKernel(XZT_C);   
    clReleaseProgram(program); 
    clReleaseCommandQueue(queue); 
    clReleaseContext(context); 
    clReleaseMemObject(CL_in5); 
    clReleaseMemObject(CL_in4); 
	clReleaseMemObject(CL_in3); 
	clReleaseMemObject(CL_in2); 
	clReleaseMemObject(CL_in1); 
	clReleaseMemObject(buffer_out); 
	clReleaseMemObject(CL_init_random); 
	clReleaseEvent(myEvent2);
	clReleaseEvent(myEvent);	
	clReleaseMemObject( CL_x );	
	clReleaseMemObject( CL_z );	
	clReleaseMemObject( CL_cc );	
	clReleaseMemObject( CL_lr );
	clReleaseMemObject( CL_ud );
	clReleaseMemObject( CL_tdis);
	clReleaseMemObject( CL_t);
	clReleaseMemObject( CL_tdisup);
	clReleaseMemObject( CL_tr);	
	clReleaseMemObject(  CL_isup);
	clReleaseMemObject(  CL_N_shift	);
	clReleaseMemObject(  CL_right__);
	clReleaseMemObject(  CL_left__	);
	clReleaseMemObject(  CL_up__);
	clReleaseMemObject(  CL_down__);
	clReleaseMemObject(  CL_uT__);
	clReleaseMemObject(  CL_dT__);
	clReleaseMemObject(  CL_lT__);
	clReleaseMemObject(  CL_rT__);
	free(host_cc);
	free(host_lr);
	free(host_ud);
	free(host_init_random);
	free(host_tr );
	free(host_tdis );
	free(host_tdisup);
	free(host_inuse);
    free(host_mw);
	return 0;
	}
	

